#include <hip/hip_runtime.h>
#include <iostream>
#include <vector>
#include <cassert>
#include <cstring>
#include <chrono>
#include <thread>
#include <cmath>

// 统一错误检查宏：封装HIP API调用错误判断，简化重复代码
// 功能：执行HIP命令，若返回非成功状态则打印错误信息并退出程序
#define HIP_CHECK(cmd)                                                                                                 \
    {                                                                                                                  \
        hipError_t error = cmd;                                                                                        \
        if (error != hipSuccess) {                                                                                     \
            std::cerr << "HIP error: " << hipGetErrorString(error) << " at line " << __LINE__ << std::endl;            \
            exit(EXIT_FAILURE);                                                                                        \
        }                                                                                                              \
    }

// 1. 全局用户数据结构：用于回调函数中传递和释放自定义数据
// 字段说明：
// - id：数据唯一标识，用于区分不同用户对象
// - value：浮点型数据，存储业务相关数值
// - name：字符数组，存储名称（固定长度32字节，避免动态内存分配）
struct UserData {
    int id;          // 数据标识
    float value;     // 浮点数值
    char name[32];   // 名称（固定长度避免动态内存）
};

// 2. 全局GPU内核定义（C++不支持函数内定义内核，需全局声明）
// 简单计算内核：将线程全局索引转为浮点值，存入设备内存
// 参数说明：
// - data：设备端内存指针，用于存储计算结果
// - size：数据总长度，用于边界检查（避免线程越界访问）
__global__ void simple_kernel(float *data, int size) {
    // 计算线程全局索引 = 块索引×块内线程数 + 块内线程索引
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    // 边界检查：仅当索引小于数据长度时执行写入（避免越界）
    if (idx < size) {
        data[idx] = idx * 1.0f;  // 索引转为浮点值存入设备内存
    }
    return;
}

// 1D纹理读取内核：适配test_texture_apis模块，仅读取1D纹理的x轴坐标数据
// 参数说明：
// - texObj：纹理对象句柄，绑定了1D数组资源
// - output：设备端内存指针，用于存储纹理读取结果
// - size：数据总长度，用于线程边界检查
__global__ void texture_read_kernel_1d(hipTextureObject_t texObj, float *output, int size) {
    // 线程索引 = 块内线程索引（仅1个线程块，直接用threadIdx.x）
    int idx = threadIdx.x;
    // 边界检查：避免线程索引超出数据长度
    if (idx >= size)
        return;

    // 1D纹理读取：使用非归一化坐标（与texDesc配置的normalizedCoords=false匹配）
    float x_coord = static_cast<float>(idx);  // 索引转为浮点坐标
    output[idx] = tex1D<float>(texObj, x_coord);  // 1D纹理专用读取函数（避免误用2D函数）

    return;
}

// Surface写入内核：适配test_surface_apis模块，支持2D坐标的Surface数据写入
// 参数说明：
// - surfObj：Surface对象句柄，绑定了2D数组资源
// - width：Surface数组的宽/高（假设为正方形，统一用width表示）
__global__ void surface_write_kernel(hipSurfaceObject_t surfObj, int width) {
    // 计算2D全局坐标：块索引×块内线程数 + 块内线程索引（支持多线程块分块执行）
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // 边界检查：确保x/y坐标不超过Surface数组的宽/高（避免越界写入）
    if (x >= width || y >= width) {
        return;
    }

    // 写入值：用x坐标作为写入数据（与主机端初始化逻辑一致，便于后续验证）
    float write_val = static_cast<float>(x);

    // Surface写入：x方向偏移 = x坐标 × 元素大小（float占4字节，匹配数组数据格式）
    surf2Dwrite<float>(write_val, surfObj, x * sizeof(float), y);
    return;
}

// 3. 用户对象销毁回调函数：全局可见，用于hipUserObjectCreate的资源释放
// 功能：当用户对象引用计数为0时，自动调用此函数释放UserData内存
// 参数说明：
// - data：用户对象绑定的自定义数据指针（指向UserData实例）
void object_destroy_callback(void *data) {
    // 空指针保护：避免传入nullptr时的非法访问
    if (data != nullptr) {
        // 类型转换：将void*转为UserData*，访问具体字段
        UserData *user_data = static_cast<UserData *>(data);
        // 打印销毁信息：便于调试跟踪用户对象的释放情况
        printf("User object destroyed: %p (id=%d, value=%.2f, name=%s)\n",
               data, user_data->id, user_data->value, user_data->name);
        delete user_data;  // 释放用户数据内存（与创建时的new配对）
    }
    return;
}

// ------------------------------ 测试图（Graph）相关API ------------------------------
// 功能：创建包含内存拷贝节点的图，验证图的实例化、执行、销毁全流程
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_graph_apis(bool should_release) {
    std::cout << "====== Testing Graph APIs ======" << std::endl;

    // 1. 创建空图：flags=0表示默认配置（无特殊属性）
    hipGraph_t graph = nullptr;
    HIP_CHECK(hipGraphCreate(&graph, 0));
    std::cout << "hipGraphCreate: SUCCESS" << std::endl;

    // 2. 准备测试数据：分配主机/设备内存并初始化主机数据
    float *d_ptr = nullptr;          // 设备端内存指针（存储拷贝目标数据）
    float *h_ptr = nullptr;          // 主机端内存指针（存储原始数据）
    const size_t data_size = 1024 * sizeof(float);  // 数据总大小（1024个float，共4KB）

    HIP_CHECK(hipMalloc(&d_ptr, data_size));  // 分配设备端内存（大小=data_size）
    h_ptr = new float[1024];                  // 分配主机端内存（1024个float）
    // 主机数据初始化：索引转为浮点值（便于后续验证拷贝结果）
    for (int i = 0; i < 1024; i++) {
        h_ptr[i] = i * 1.0f;
    }

    // 3. 向图中添加内存拷贝节点：实现主机→设备的数据拷贝
    hipGraphNode_t memcpy_node;  // 图节点句柄（存储内存拷贝节点）
    HIP_CHECK(hipGraphAddMemcpyNode1D(
        &memcpy_node,    // 输出参数：内存拷贝节点句柄
        graph,           // 所属图：将节点添加到当前创建的graph中
        nullptr,         // 依赖节点列表：无依赖（当前节点为第一个节点）
        0,               // 依赖节点数量：0（无依赖）
        d_ptr,           // 目标地址：设备端内存指针
        h_ptr,           // 源地址：主机端内存指针
        data_size,       // 拷贝字节数：data_size（4KB）
        hipMemcpyHostToDevice  // 拷贝方向：主机→设备
    ));

    // 4. 实例化图并执行：将图转为可执行实例，在默认流中运行
    hipGraphExec_t graph_exec = nullptr;  // 图执行实例句柄
    HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0));
    HIP_CHECK(hipGraphLaunch(graph_exec, 0));  // 在默认流（0）中执行图
    HIP_CHECK(hipStreamSynchronize(0));         // 等待默认流执行完成（确保图运行结束）

    // 5. 验证图执行结果：设备→主机拷贝数据，对比原始数据与拷贝结果
    float *h_verify = new float[1024];  // 主机端内存指针（存储验证数据）
    HIP_CHECK(hipMemcpy(h_verify, d_ptr, data_size, hipMemcpyDeviceToHost));

    bool success = true;
    // 浮点数据对比：允许1e-5的误差（避免浮点精度问题导致的误判）
    for (int i = 0; i < 1024; i++) {
        if (fabs(h_verify[i] - h_ptr[i]) > 1e-5) {
            success = false;
        }
    }
    std::cout << "Graph execution verified: " << (success ? "SUCCESS" : "FAILED") << std::endl;

    // 6. 清理临时资源：释放设备内存和主机内存（无论是否泄漏图，临时资源都需释放）
    HIP_CHECK(hipFree(d_ptr));        // 释放设备端内存（与hipMalloc配对）
    delete[] h_ptr;                   // 释放主机端原始数据内存（与new配对）
    delete[] h_verify;                // 释放主机端验证数据内存（与new配对）

    // 7. 条件释放图资源：仅当should_release=true时释放（避免泄漏时的冗余打印）
    if (should_release) {
        HIP_CHECK(hipGraphExecDestroy(graph_exec));  // 销毁图执行实例（与hipGraphInstantiate配对）
        HIP_CHECK(hipGraphDestroy(graph));           // 销毁图（与hipGraphCreate配对）
        std::cout << "hipGraphDestroy + hipGraphExecDestroy: SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 测试数组（Array）相关API ------------------------------
// 功能：验证1D数组、3D数组、Mipmap数组的创建、数据拷贝、纹理绑定、销毁全流程
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_array_apis(bool should_release) {
    std::cout << "====== Testing Array APIs ======" << std::endl;

    // ------------------------------ 1D数组测试 ------------------------------
    // 1D数组描述符：初始化清零（避免未初始化字段的干扰）
    HIP_ARRAY_DESCRIPTOR arrayDesc = {};
    arrayDesc.Width = 256;                // 1D数组宽度：256个元素
    arrayDesc.Height = 0;                 // 1D数组高度：固定为0（区分2D数组）
    arrayDesc.Format = HIP_AD_FORMAT_FLOAT;// 数据格式：单精度浮点（32位）
    arrayDesc.NumChannels = 1;            // 通道数：1（灰度数据，无RGB/A通道）

    hipArray_t array = nullptr;
    HIP_CHECK(hipArrayCreate(&array, &arrayDesc));  // 创建1D数组
    std::cout << "hipArrayCreate (1D): SUCCESS" << std::endl;

    // 主机→1D数组拷贝：使用hipMemcpy2DToArray（适配数组的存储格式）
    float *h_data = new float[256];  // 主机端内存：存储1D数组的原始数据
    // 主机数据初始化：索引转为浮点值
    for (int i = 0; i < 256; i++) {
        h_data[i] = i * 1.0f;
    }

    HIP_CHECK(hipMemcpy2DToArray(
        array,                          // 目标：1D数组
        0, 0,                           // 数组内偏移：x=0，y=0（1D数组y偏移无效）
        h_data,                         // 源：主机端内存指针
        256 * sizeof(float),            // 源数据行间距：1D数组=宽度×元素大小（无换行）
        256 * sizeof(float),            // 每行拷贝字节数：与宽度匹配（256个float）
        1,                              // 拷贝行数：1D数组仅1行
        hipMemcpyHostToDevice           // 拷贝方向：主机→设备
    ));
    std::cout << "1D Array data copied (Host→Device): SUCCESS" << std::endl;
    delete[] h_data;  // 释放主机端原始数据内存（拷贝完成后无需保留）

    // ------------------------------ 3D数组测试 ------------------------------
    // 3D数组描述符：初始化清零
    HIP_ARRAY3D_DESCRIPTOR array3DDesc = {};
    array3DDesc.Width = 64;                   // 宽度：64个元素
    array3DDesc.Height = 64;                  // 高度：64个元素
    array3DDesc.Depth = 64;                   // 深度：64个元素
    array3DDesc.Format = HIP_AD_FORMAT_FLOAT; // 数据格式：单精度浮点
    array3DDesc.NumChannels = 1;              // 通道数：1

    hipArray_t array3D = nullptr;
    HIP_CHECK(hipArray3DCreate(&array3D, &array3DDesc));  // 创建3D数组
    std::cout << "hipArray3DCreate: SUCCESS" << std::endl;

    // 3D数组绑定纹理：验证数组与纹理对象的兼容性
    hipResourceDesc resDesc = {};
    resDesc.resType = hipResourceTypeArray;    // 资源类型：数组（绑定3D数组）
    resDesc.res.array.array = array3D;         // 绑定的3D数组句柄

    hipTextureDesc texDesc = {};
    texDesc.addressMode[0] = hipAddressModeWrap;  // x轴地址模式：循环（超出范围后_wrap_）
    texDesc.filterMode = hipFilterModeLinear;     // 滤波模式：线性插值（平滑过渡）
    texDesc.readMode = hipReadModeElementType;    // 读取模式：原始数据类型（不做格式转换）

    hipTextureObject_t texObj = nullptr;
    HIP_CHECK(hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr));  // 创建纹理对象
    std::cout << "3D Array bound to Texture Object: SUCCESS" << std::endl;
    HIP_CHECK(hipDestroyTextureObject(texObj));  // 销毁纹理对象（仅验证绑定，无需保留）

    // ------------------------------ Mipmap数组测试 ------------------------------
    // Mipmap数组描述符：2D基础层+多级缩放（Depth=1表示2D Mipmap）
    HIP_ARRAY3D_DESCRIPTOR mipmapDesc = {};
    mipmapDesc.Width = 128;                  // 基础层宽度：128个元素
    mipmapDesc.Height = 128;                 // 基础层高度：128个元素
    mipmapDesc.Depth = 1;                    // 深度：1（2D Mipmap，无深度维度）
    mipmapDesc.Format = HIP_AD_FORMAT_FLOAT; // 数据格式：单精度浮点
    mipmapDesc.NumChannels = 1;              // 通道数：1

    hipMipmappedArray_t mipmapArray = nullptr;
    // 创建4级Mipmap数组：基础层128x128 → 64x64 → 32x32 → 16x16（共4级）
    HIP_CHECK(hipMipmappedArrayCreate(&mipmapArray, &mipmapDesc, 4));
    std::cout << "hipMipmappedArrayCreate (4 levels): SUCCESS" << std::endl;

    // 访问各级Mipmap：验证多级索引的有效性（从0级到3级）
    for (unsigned int level = 0; level < 4; level++) {
        hipArray_t levelArray = nullptr;
        HIP_CHECK(hipMipmappedArrayGetLevel(&levelArray, mipmapArray, level));
        std::cout << "Mipmap level " << level << " accessed: SUCCESS" << std::endl;
    }

    // ------------------------------ 资源释放 ------------------------------
    // 仅当should_release=true时释放数组资源（避免泄漏时的冗余打印）
    if (should_release) {
        HIP_CHECK(hipArrayDestroy(array));         // 释放1D数组（与hipArrayCreate配对）
        std::cout << "hipArrayDestroy (1D): SUCCESS" << std::endl;

        HIP_CHECK(hipArrayDestroy(array3D));       // 释放3D数组（与hipArray3DCreate配对）
        std::cout << "hipArrayDestroy (3D): SUCCESS" << std::endl;

        HIP_CHECK(hipMipmappedArrayDestroy(mipmapArray));  // 释放Mipmap数组（与hipMipmappedArrayCreate配对）
        std::cout << "hipMipmappedArrayDestroy: SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 测试上下文（Context）相关API ------------------------------
// 功能：验证显式上下文的创建、切换、设备计算、销毁全流程
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_context_apis(bool should_release) {
    std::cout << "====== Testing Context APIs ======" << std::endl;

    // 1. 显式创建上下文：绑定设备0，flags=0表示默认配置（无特殊属性）
    hipCtx_t ctx = nullptr;
    int deviceId = 0;
    HIP_CHECK(hipCtxCreate(&ctx, 0, deviceId));
    std::cout << "hipCtxCreate: SUCCESS (ctx=" << static_cast<void *>(ctx) << ")" << std::endl;

    // 2. 切换到显式创建的上下文：确保后续HIP操作在当前上下文执行
    HIP_CHECK(hipCtxSetCurrent(ctx));
    std::cout << "hipCtxSetCurrent: SUCCESS (switched to explicit context)" << std::endl;

    // 3. 验证上下文绑定的设备信息：获取设备属性，确认上下文与设备的绑定关系
    hipDeviceProp_t prop;
    HIP_CHECK(hipGetDeviceProperties(&prop, deviceId));
    std::cout << "Using device under context: " << prop.name << std::endl;

    // 4. 在当前上下文执行计算：验证上下文的有效性（设备内存分配→数据拷贝→内核执行→结果验证）
    float *d_data = nullptr;          // 设备端内存指针（存储计算数据）
    float *h_data = new float[1024];  // 主机端输入数据指针
    float *h_result = new float[1024];// 主机端结果数据指针

    // 分配设备端内存（大小=1024个float，共4KB）
    HIP_CHECK(hipMalloc(&d_data, 1024 * sizeof(float)));
    // 主机数据初始化：索引转为浮点值
    for (int i = 0; i < 1024; i++) {
        h_data[i] = i * 1.0f;
    }

    // 主机→设备拷贝：将原始数据传入设备内存
    HIP_CHECK(hipMemcpy(d_data, h_data, 1024 * sizeof(float), hipMemcpyHostToDevice));
    // 执行内核：32个线程块 × 32个线程/块 = 1024个线程（覆盖所有数据）
    simple_kernel<<<32, 32>>>(d_data, 1024);
    HIP_CHECK(hipDeviceSynchronize());  // 等待内核执行完成（确保计算结束）

    // 设备→主机拷贝：将计算结果传回主机内存
    HIP_CHECK(hipMemcpy(h_result, d_data, 1024 * sizeof(float), hipMemcpyDeviceToHost));

    // 验证计算结果：对比原始数据与计算结果（内核仅将索引转为浮点，结果应与原始数据一致）
    bool success = true;
    for (int i = 0; i < 1024; i++) {
        if (fabs(h_result[i] - h_data[i]) > 1e-5) {
            success = false;
        }
    }
    std::cout << "Context computation verified: " << (success ? "SUCCESS" : "FAILED") << std::endl;

    // 5. 清理计算资源：无论是否泄漏上下文，临时资源都需释放
    HIP_CHECK(hipFree(d_data));        // 释放设备端内存（与hipMalloc配对）
    delete[] h_data;                   // 释放主机端输入数据内存（与new配对）
    delete[] h_result;                 // 释放主机端结果数据内存（与new配对）

    // 6. 条件释放上下文：仅当should_release=true时释放（避免泄漏时的冗余打印）
    if (should_release) {
        HIP_CHECK(hipCtxDestroy(ctx));  // 销毁显式上下文（与hipCtxCreate配对）
        std::cout << "hipCtxDestroy (explicit context): SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 测试纹理（Texture）相关API ------------------------------
// 功能：基于1D数组验证纹理对象的创建、数据读取、销毁全流程（对齐正常用例）
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_texture_apis(bool should_release) {
    std::cout << "====== Testing Texture APIs ======" << std::endl;

    // 1. 配置1D数组参数：与纹理对象的1D类型匹配（避免维度不兼容）
    const int width = 256;  // 1D数组宽度：256个元素（无高度）
    // 通道描述符：明确32位单精度浮点，单通道（避免纹理格式解析错误）
    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
    hipArray *array = nullptr;

    // 分配1D数组：仅指定宽度（无高度参数，明确为1D数组）
    HIP_CHECK(hipMallocArray(&array, &channelDesc, width));
    std::cout << "hipMallocArray (1D for Texture): SUCCESS (array=" << static_cast<void *>(array) << ")" << std::endl;

    // 2. 主机→1D数组数据拷贝：行数=1（匹配1D数组的维度）
    float *h_texture_data = new float[width];  // 主机端原始数据指针
    // 主机数据初始化：非归一化数据（与纹理的normalizedCoords=false匹配）
    for (int i = 0; i < width; i++) {
        h_texture_data[i] = static_cast<float>(i);
    }

    HIP_CHECK(hipMemcpy2DToArray(
        array,                          // 目标：1D数组
        0, 0,                           // 偏移：x=0，y=0（1D数组y偏移无效）
        h_texture_data,                 // 源：主机端原始数据
        width * sizeof(float),          // 源行间距：1D数组=宽度×元素大小（无换行）
        width * sizeof(float),          // 每行拷贝字节数：与宽度匹配
        1,                              // 拷贝行数：1D数组仅1行
        hipMemcpyHostToDevice           // 拷贝方向：主机→设备
    ));
    std::cout << "1D Texture data copied (Host→Device): SUCCESS" << std::endl;
    delete[] h_texture_data;  // 释放主机端原始数据（拷贝完成后无需保留）

    // 3. 配置纹理资源描述符：仅设置1D相关字段（避免冗余字段干扰）
    hipResourceDesc resDesc = {};
    memset(&resDesc, 0, sizeof(resDesc));  // 显式清零（避免未初始化字段的随机值）
    resDesc.resType = hipResourceTypeArray;    // 资源类型：数组（绑定1D数组）
    resDesc.res.array.array = array;           // 绑定的1D数组句柄

    // 4. 配置纹理描述符：1D纹理专用配置（删除2D相关字段）
    hipTextureDesc texDesc = {};
    memset(&texDesc, 0, sizeof(texDesc));  // 显式清零（符合HIP最佳实践）
    texDesc.addressMode[0] = hipAddressModeClamp;  // x轴地址模式：截断（超出范围取边界值）
    texDesc.filterMode = hipFilterModePoint;       // 滤波模式：点采样（取最近像素，兼容性最好）
    texDesc.readMode = hipReadModeElementType;      // 读取模式：原始数据类型（不做格式转换）
    texDesc.normalizedCoords = false;              // 禁用归一化坐标（与主机数据格式匹配）

    // 5. 创建纹理对象并等待初始化：确保纹理资源完全就绪（避免读取时未初始化）
    hipTextureObject_t texObj = 0;  // 初始化为0（符合HIP API的默认值习惯）
    HIP_CHECK(hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr));
    HIP_CHECK(hipDeviceSynchronize());  // 等待纹理对象初始化完成
    std::cout << "hipCreateTextureObject (1D): SUCCESS (texObj=" << static_cast<void *>(texObj) << ")" << std::endl;

    // 6. 分配纹理读取结果的设备内存：存储纹理内核的读取结果
    float *d_output = nullptr;
    HIP_CHECK(hipMalloc(&d_output, width * sizeof(float)));
    std::cout << "hipMalloc for Texture output: SUCCESS (d_output=" << static_cast<void *>(d_output) << ")" << std::endl;

    // 7. 执行1D纹理读取内核：线程配置与1D数组匹配（1个线程块 × width个线程）
    std::cout << "Launching 1D texture_read_kernel_1d..." << std::endl;
    texture_read_kernel_1d<<<1, width>>>(texObj, d_output, width);
    HIP_CHECK(hipGetLastError());     // 检查内核启动参数错误（如线程数超限）
    HIP_CHECK(hipDeviceSynchronize());// 等待内核执行完成（捕获执行阶段错误）

    // 8. 验证纹理读取结果：设备→主机拷贝，对比预期值与实际结果
    float *h_output = new float[width];  // 主机端结果数据指针
    HIP_CHECK(hipMemcpy(h_output, d_output, width * sizeof(float), hipMemcpyDeviceToHost));

    bool tex_success = true;
    for (int i = 0; i < width; i++) {
        float expected = static_cast<float>(i);  // 预期值=主机原始数据（索引转为浮点）
        if (fabs(h_output[i] - expected) > 1e-5) {
            std::cerr << "Mismatch at index " << i << ": expected " << expected << ", got " << h_output[i] << std::endl;
            tex_success = false;
        }
    }
    std::cout << "1D Texture read verified: " << (tex_success ? "SUCCESS" : "FAILED") << std::endl;

    // 9. 清理临时资源：无论是否泄漏纹理，临时内存都需释放
    HIP_CHECK(hipFree(d_output));    // 释放设备端结果内存（与hipMalloc配对）
    delete[] h_output;               // 释放主机端结果内存（与new配对）

    // 10. 条件释放纹理资源：仅当should_release=true时释放（避免泄漏时的冗余打印）
    if (should_release) {
        HIP_CHECK(hipDestroyTextureObject(texObj));  // 销毁纹理对象（与hipCreateTextureObject配对）
        std::cout << "hipDestroyTextureObject: SUCCESS" << std::endl;

        HIP_CHECK(hipFreeArray(array));              // 释放1D数组（与hipMallocArray配对）
        std::cout << "hipFreeArray (1D): SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 测试用户对象（User Object）API ------------------------------
// 功能：验证用户对象的创建、释放流程（适配旧版HIP参数）
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void TestFunction(void *userData) {
    // 空回调函数：用于hipUserObjectCreate的参数（无实际业务逻辑，仅占位）
    int a = 0;
    return;
}

void test_user_object_apis(bool should_release) {
    std::cout << "====== Testing User Object APIs ====== " << std::endl;

    hipUserObject_t object_out = nullptr;  // 用户对象句柄（输出参数）
    // 自定义数据指针：示例值12345（实际场景可指向UserData等自定义结构）
    void *ptr = reinterpret_cast<void *>(12345);
    unsigned int initialRefcount = 1;      // 初始引用计数：1（创建时默认引用1次）
    // 标志：无析构同步（适配旧版HIP支持的标志，避免使用未定义标志）
    unsigned int flags = hipUserObjectNoDestructorSync;

    // 创建用户对象：绑定自定义数据和回调函数
    int result = hipUserObjectCreate(&object_out, ptr, TestFunction, initialRefcount, flags);

    if (result == hipSuccess && object_out != nullptr) {
        std::cout << "hipUserObjectCreate: SUCCESS" << std::endl;

        // 条件释放用户对象：仅当should_release=true时释放（引用计数需与创建时匹配）
        if (should_release) {
            hipError_t releaseResult = hipUserObjectRelease(object_out, initialRefcount);
            if (releaseResult == hipSuccess) {
                std::cout << "hipUserObjectRelease: SUCCESS" << std::endl;
            } else {
                std::cerr << "hipUserObjectRelease: FAILED (" << hipGetErrorString(releaseResult) << ")" << std::endl;
            }
        }
    }

    return;
}

// ------------------------------ 测试内存管理（Memory）API ------------------------------
// 功能：验证通用内存创建（hipMemCreate）和内存池（hipMemPool）的流程
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_memory_apis(bool should_release) {
    std::cout << "====== Testing Memory APIs ======" << std::endl;

    // ------------------------------ 测试hipMemCreate（通用内存分配） ------------------------------
    hipMemGenericAllocationHandle_t handle = nullptr;  // 通用内存句柄
    size_t size = 4 * 1024;                            // 分配大小：4KB（适合小内存测试）
    hipMemAllocationProp Prop = {};                     // 内存分配属性（初始化清零）
    Prop.type = hipMemAllocationTypePinned;             // 内存类型：固定内存（Pinned，主机可直接访问）
    Prop.location.id = 0;                               // 绑定设备ID：0（与当前设备一致）
    Prop.location.type = hipMemLocationTypeDevice;       // 位置类型：设备（内存物理位置在设备端）
    Prop.requestedHandleType = hipMemHandleTypeNone;     // 句柄类型：无（不使用外部句柄）
    Prop.win32HandleMetaData = nullptr;                 // Windows句柄元数据：未使用（Linux环境）
    Prop.allocFlags.compressionType = 0;                // 压缩类型：无（不启用内存压缩）
    Prop.allocFlags.gpuDirectRDMACapable = 0;           // GPU Direct RDMA：禁用（简化测试）
    Prop.allocFlags.usage = 0;                          // 用途：默认（无特殊用途）

    unsigned long long flags = 0;  // 分配标志：默认（无特殊属性）
    size_t granularity = 0;        // 分配粒度：后续通过API获取

    // 获取推荐的分配粒度：确保内存分配符合设备要求（避免对齐错误）
    HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &Prop, hipMemAllocationGranularityRecommended));
    // 创建通用内存：按指定属性和大小分配
    HIP_CHECK(hipMemCreate(&handle, size, &Prop, flags));
    std::cout << "hipMemCreate (4KB Pinned Memory): SUCCESS" << std::endl;

    // ------------------------------ 测试内存池（hipMemPool） ------------------------------
    hipMemPool_t memPool = nullptr;  // 内存池句柄
    hipMemPoolProps poolProps = {};  // 内存池属性（初始化清零）
    poolProps.allocType = hipMemAllocationTypePinned;       // 内存类型：固定内存（与通用内存一致）
    poolProps.location.type = hipMemLocationTypeDevice;     // 位置类型：设备
    poolProps.location.id = 0;                               // 绑定设备ID：0

    // 创建内存池：按指定属性初始化
    hipError_t poolCreateResult = hipMemPoolCreate(&memPool, &poolProps);
    if (poolCreateResult == hipSuccess) {
        std::cout << "hipMemPoolCreate: SUCCESS" << std::endl;

        // 从内存池分配内存：异步分配（默认流），大小4KB
        void *poolAllocation = nullptr;
        HIP_CHECK(hipMallocFromPoolAsync(
            &poolAllocation,    // 输出：分配的内存指针
            4096,               // 分配大小：4KB
            memPool,            // 所属内存池：当前创建的memPool
            0                   // 流：默认流（0）
        ));

        // 执行简单内核：验证内存池分配的内存可正常访问
        simple_kernel<<<1, 64>>>((float *)poolAllocation, 64);  // 1块×64线程，处理64个数据
        HIP_CHECK(hipStreamSynchronize(0));  // 等待内核执行完成

        // 验证计算结果：设备→主机拷贝，对比预期值
        float *h_pool_result = new float[64];  // 主机端结果指针
        HIP_CHECK(hipMemcpy(h_pool_result, poolAllocation, 64 * sizeof(float), hipMemcpyDeviceToHost));

        bool pool_success = true;
        for (int i = 0; i < 64; i++) {
            if (fabs(h_pool_result[i] - i * 1.0f) > 1e-5) {
                pool_success = false;
            }
        }
        std::cout << "Memory Pool computation verified: " << (pool_success ? "SUCCESS" : "FAILED") << std::endl;

        // 清理内存池分配的临时资源
        HIP_CHECK(hipFree(poolAllocation));  // 释放内存池分配的内存（与hipMallocFromPoolAsync配对）
        delete[] h_pool_result;              // 释放主机端结果内存（与new配对）

        // 条件释放内存池和通用内存：仅当should_release=true时释放
        if (should_release) {
            HIP_CHECK(hipMemRelease(handle));          // 释放通用内存（与hipMemCreate配对）
            std::cout << "hipMemRelease (Generic Memory): SUCCESS" << std::endl;

            HIP_CHECK(hipMemPoolDestroy(memPool));     // 释放内存池（与hipMemPoolCreate配对）
            std::cout << "hipMemPoolDestroy: SUCCESS" << std::endl;
        }
    } else {
        // 内存池创建失败：打印提示（可能是HIP版本不支持）
        std::cout << "hipMemPoolCreate: NOT SUPPORTED (error code: " << poolCreateResult << ")" << std::endl;
        // 仅释放通用内存（无论是否泄漏，创建成功的资源需释放）
        if (should_release) {
            HIP_CHECK(hipMemRelease(handle));
            std::cout << "hipMemRelease (Generic Memory): SUCCESS" << std::endl;
        }
    }

    return;
}

// ------------------------------ 测试Surface相关API（基础版） ------------------------------
// 功能：基于2D数组验证Surface对象的创建、数据拷贝、销毁全流程（无内核执行）
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_surface_apis(bool should_release) {
    std::cout << "====== Testing Surface APIs ======" << std::endl;

    // 1. 配置2D数组参数：与Surface对象的2D类型匹配
    int width = 100;   // 2D数组宽度：100个元素
    int height = 50;   // 2D数组高度：50个元素
    // 主机端2D数组：存储Surface的原始数据（栈上分配，避免动态内存）
    float surfaceData[width][height];
    // 主机数据初始化：所有元素设为1.0f（简化后续验证逻辑）
    for (int i = 0; i < width; i++) {
        for (int j = 0; j < height; j++) {
            surfaceData[i][j] = 1.0f;
        }
    }

    // 2. 分配设备2D数组：用于绑定Surface对象（存储设备端Surface数据）
    hipArray* d_surfaceArray = nullptr;
    // 通道描述符：单精度浮点，单通道（适配Surface的默认格式）
    hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
    HIP_CHECK(hipMallocArray(&d_surfaceArray, &channelDesc, width, height));
    std::cout << "hipMallocArray (2D for Surface): SUCCESS" << std::endl;

    // 3. 主机→设备2D数组拷贝：初始化Surface的设备端数据
    HIP_CHECK(hipMemcpy2DToArray(
        d_surfaceArray,                  // 目标：2D数组
        0, 0,                           // 偏移：x=0，y=0（从数组起始位置拷贝）
        surfaceData,                     // 源：主机端2D数组
        width * sizeof(float),            // 源行间距：每行字节数（width个float）
        width * sizeof(float),            // 每行拷贝字节数：与源行间距一致
        height,                          // 拷贝行数：数组高度（50行）
        hipMemcpyHostToDevice             // 拷贝方向：主机→设备
    ));
    std::cout << "Surface data copied (Host→Device): SUCCESS" << std::endl;

    // 4. 配置Surface资源描述符：绑定2D数组资源
    hipResourceDesc resDesc = {};
    memset(&resDesc, 0, sizeof(resDesc));  // 显式清零（避免未初始化字段）
    resDesc.resType = hipResourceTypeArray;    // 资源类型：数组（绑定2D数组）
    resDesc.res.array.array = d_surfaceArray;  // 绑定的2D数组句柄

    // 5. 创建Surface对象：按资源描述符初始化
    hipSurfaceObject_t surfaceObject = nullptr;
    HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
    std::cout << "hipCreateSurfaceObject: SUCCESS" << std::endl;

    // 6. 条件释放Surface资源：仅当should_release=true时释放
    if (should_release) {
        HIP_CHECK(hipDestroySurfaceObject(surfaceObject));  // 销毁Surface对象（与hipCreateSurfaceObject配对）
        std::cout << "hipDestroySurfaceObject: SUCCESS" << std::endl;

        HIP_CHECK(hipFreeArray(d_surfaceArray));            // 释放2D数组（与hipMallocArray配对）
        std::cout << "hipFreeArray (2D): SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 测试Surface相关API（兼容旧版函数名） ------------------------------
// 功能：使用旧版hipSurfObjectCreate函数，验证Surface创建/销毁（适配特定HIP版本）
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_surface_apis2(bool should_release) {
    std::cout << "====== Testing Surface APIs (Old: hipSurfObjectCreate) ======" << std::endl;

    int width = 100;   // 2D数组宽度：100个元素
    int height = 50;   // 2D数组高度：50个元素
    // 主机端2D数组：存储原始数据
    float surfaceData[width][height];
    for (int i = 0; i < width; i++) {
        for (int j = 0; j < height; j++) {
            surfaceData[i][j] = 1.0f;
        }
    }

    // 分配设备2D数组
    hipArray* d_surfaceArray = nullptr;
    hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
    HIP_CHECK(hipMallocArray(&d_surfaceArray, &channelDesc, width, height));

    // 主机→设备数据拷贝
    HIP_CHECK(hipMemcpy2DToArray(
        d_surfaceArray,                  // 目标：2D数组
        0, 0,                           // 偏移：x=0，y=0
        surfaceData,                     // 源：主机端数据
        width * sizeof(float),            // 源行间距：每行字节数
        width * sizeof(float),            // 每行拷贝字节数
        height,                          // 拷贝行数：数组高度
        hipMemcpyHostToDevice             // 拷贝方向：主机→设备
    ));

    // 配置旧版Surface资源描述符（HIP_RESOURCE_DESC，适配旧版API）
    HIP_RESOURCE_DESC pResDesc = {};
    memset(&pResDesc, 0, sizeof(HIP_RESOURCE_DESC));
    pResDesc.resType = HIP_RESOURCE_TYPE_ARRAY;    // 资源类型：数组
    pResDesc.res.array.hArray = d_surfaceArray;    // 绑定的2D数组句柄
    pResDesc.flags = 0;                            // 标志：默认

    // 旧版函数创建Surface对象：适配不支持hipCreateSurfaceObject的HIP版本
    hipSurfaceObject_t pSurfObject = nullptr;
    HIP_CHECK(hipSurfObjectCreate(&pSurfObject, &pResDesc));

    // 条件释放：仅当should_release=true时释放
    if (should_release) {
        HIP_CHECK(hipSurfObjectDestroy(pSurfObject));  // 销毁旧版Surface对象（与hipSurfObjectCreate配对）
        HIP_CHECK(hipFreeArray(d_surfaceArray));        // 释放2D数组（与hipMallocArray配对）
        std::cout << "hipSurfObjectDestroy + hipFreeArray: SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 测试弃用纹理引用（Texture Reference）API ------------------------------
// 功能：验证旧版纹理引用的创建、绑定、销毁流程（忽略弃用警告）
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_texref_apis(bool should_release) {
    std::cout << "====== Testing Deprecated Texture Reference APIs ======" << std::endl;

    // 忽略弃用API的编译警告：旧版HIP需兼容，暂不替换为纹理对象
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wdeprecated-declarations"

    textureReference *texRef = nullptr;
    // 创建纹理引用：旧版API，仅支持flags=0（默认配置）
    hipError_t texRefResult = hipTexRefCreate(&texRef);
    if (texRefResult == hipSuccess) {
        std::cout << "hipTexRefCreate: SUCCESS (texRef=" << static_cast<void *>(texRef) << ")" << std::endl;

        // 绑定1D数组到纹理引用：创建适配的1D数组
        HIP_ARRAY_DESCRIPTOR texRefArrayDesc = {};
        texRefArrayDesc.Width = 128;                // 1D数组宽度：128个元素
        texRefArrayDesc.Height = 0;                 // 1D数组高度：0（区分2D）
        texRefArrayDesc.Format = HIP_AD_FORMAT_FLOAT;// 数据格式：单精度浮点
        texRefArrayDesc.NumChannels = 1;            // 通道数：1

        hipArray_t texRefArray = nullptr;
        HIP_CHECK(hipArrayCreate(&texRefArray, &texRefArrayDesc));

        // 配置纹理引用参数：覆盖格式、地址模式、滤波模式
        HIP_CHECK(hipTexRefSetArray(texRef, texRefArray, HIP_TRSA_OVERRIDE_FORMAT));  // 绑定数组并覆盖格式
        HIP_CHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeWrap));  // x轴地址模式：循环
        HIP_CHECK(hipTexRefSetFilterMode(texRef, hipFilterModeLinear));     // 滤波模式：线性插值
        std::cout << "Texture Reference configured and bound to 1D Array: SUCCESS" << std::endl;

        // 条件释放：仅当should_release=true时释放
        if (should_release) {
            HIP_CHECK(hipTexRefDestroy(texRef));  // 销毁纹理引用（与hipTexRefCreate配对）
            std::cout << "hipTexRefDestroy: SUCCESS" << std::endl;

            HIP_CHECK(hipArrayDestroy(texRefArray));  // 释放绑定的1D数组（与hipArrayCreate配对）
            std::cout << "hipArrayDestroy (bound to TexRef): SUCCESS" << std::endl;
        }
    } else {
        // 纹理引用创建失败：打印错误信息（便于调试）
        std::cout << "hipTexRefCreate: FAILED (error: " << hipGetErrorString(texRefResult) << ")" << std::endl;
    }

    // 恢复编译警告配置：避免影响其他代码
#pragma GCC diagnostic pop

    return;
}

// ------------------------------ 测试实验性API（hipGreenCtx） ------------------------------
// 功能：验证实验性Green Context的资源拆分、创建、销毁全流程
// 参数说明：
// - should_release：是否释放资源（true=正常释放，false=故意泄漏）
void test_hipGreenCtxCreate_apis(bool should_release) {
    std::cout << "====== Testing Experimental APIs (hipGreenCtx) ======" << std::endl;

    hipDevResource input;                  // 输入设备资源（SM资源）
    unsigned int nbGroups = 10;           // 资源拆分组数：10组
    unsigned int minCount = 10;           // 每组最小资源数：10
    hipDevResource out_resources[nbGroups];// 输出拆分后的资源数组（10组）
    hipDevResource remain_resource;        // 拆分后剩余的资源
    const int N = 10240;                  // 数据长度：10240个int（用于内核测试）
    int* d_data = nullptr;                // 设备端内存指针（存储内核计算数据）

    // 分配设备端内存：用于内核计算验证
    HIP_CHECK(hipMalloc((void**)&d_data, N * sizeof(int)));

    // 获取设备SM资源：用于Green Context的资源拆分
    HIP_CHECK(hipDeviceGetDevResource(0, &input, HIP_DEV_RESOURCE_TYPE_SM));
    // 拆分SM资源：按组数和最小资源数拆分
    HIP_CHECK(hipDevSmResourceSplitByCount(
        out_resources,    // 输出：拆分后的资源数组
        &nbGroups,        // 输入输出：实际拆分组数（与输入一致）
        &input,           // 输入：原始SM资源
        &remain_resource, // 输出：拆分后剩余资源
        0,                // 标志：默认
        minCount          // 每组最小资源数
    ));

    // 初始化Green Context相关资源：描述符、句柄、流
    hipDevResourceDesc ctxDesc[nbGroups];  // Green Context描述符数组（10个）
    hipGreenCtx greenCtx[nbGroups];        // Green Context句柄数组（10个）
    hipStream_t stream[nbGroups];          // 流数组（10个，绑定到Green Context）
    for (int i = 0; i < nbGroups; i++) {
        // 生成Green Context描述符：基于拆分后的资源
        HIP_CHECK(hipDevResourceGenerateDesc(&ctxDesc[i], &out_resources[i], 1));
        // 创建Green Context：按描述符初始化
        HIP_CHECK(hipGreenCtxCreate(&greenCtx[i], ctxDesc[i], 0, HIP_GREEN_CTX_DEFAULT_STREAM));
        // 创建Green Context绑定的流：非阻塞模式
        HIP_CHECK(hipGreenCtxStreamCreate(&stream[i], greenCtx[i], hipStreamNonBlocking, 0));
    }

    // 获取Green Context对应的常规上下文，并切换到该上下文
    hipCtx_t ctx;
    HIP_CHECK(hipCtxFromGreenCtx(&ctx, greenCtx[0]));
    HIP_CHECK(hipCtxSetCurrent(ctx));

    // 创建事件：用于测量内核执行时间
    hipEvent_t startevent, stopEvent;
    HIP_CHECK(hipEventCreate(&startevent));
    HIP_CHECK(hipEventCreate(&stopEvent));

    // 记录开始事件，执行内核，记录结束事件
    HIP_CHECK(hipGreenCtxRecordEvent(greenCtx[0], startevent));
    // 执行内核：修复笔误（simpleKernel→simple_kernel，与全局内核名一致）
    simple_kernel<<<N / 256, 256, 0, stream[0]>>>((float*)d_data, N);
    HIP_CHECK(hipGreenCtxRecordEvent(greenCtx[0], stopEvent));

    // 等待事件完成，计算并打印内核执行时间
    HIP_CHECK(hipEventSynchronize(stopEvent));
    float elapsedTime;
    HIP_CHECK(hipEventElapsedTime(&elapsedTime, startevent, stopEvent));
    std::cout << "Kernel execution time under Green Context: " << elapsedTime << " ms" << std::endl;

    // 清理临时资源：无论是否泄漏Green Context，临时内存和事件都需释放
    HIP_CHECK(hipFree(d_data));                // 释放设备端内存（与hipMalloc配对）
    HIP_CHECK(hipEventDestroy(startevent));    // 销毁开始事件（与hipEventCreate配对）
    HIP_CHECK(hipEventDestroy(stopEvent));     // 销毁结束事件（与hipEventCreate配对）

    // 条件释放Green Context资源：仅当should_release=true时释放
    if (should_release) {
        for (int i = 0; i < nbGroups; i++) {
            HIP_CHECK(hipStreamDestroy(stream[i]));  // 销毁流（与hipGreenCtxStreamCreate配对）
            HIP_CHECK(hipGreenCtxDestroy(greenCtx[i]));  // 销毁Green Context（与hipGreenCtxCreate配对）
        }
        std::cout << "hipGreenCtxDestroy + hipStreamDestroy: SUCCESS" << std::endl;
    }

    return;
}

// ------------------------------ 主函数 ------------------------------
// 功能：解析命令行参数，初始化HIP环境，执行所有测试模块，输出测试总结
int main(int argc, char **argv) {
    std::cout << "=== Starting HIP Memory Leak Detection Test ===" << std::endl;
    std::cout << "Test Version: All allocations are actually used" << std::endl;

    // 解析命令行参数：--leak表示故意泄漏资源，否则正常释放
    bool release_resources = true;
    if (argc > 1 && std::string(argv[1]) == "--leak") {
        release_resources = false;
        std::cout << "Test Mode: INTENTIONAL LEAK (resources not released)" << std::endl;
    } else {
        std::cout << "Test Mode: CLEAN RELEASE (all resources released)" << std::endl;
    }

    // 1. 初始化HIP环境：必须在所有HIP API调用前执行
    hipError_t initResult = hipInit(0);
    if (initResult != hipSuccess) {
        std::cerr << "Failed to initialize HIP: " << hipGetErrorString(initResult) << std::endl;
        return -1;
    }

    // 2. 检查并设置设备：确保有可用HIP设备，默认使用第0号设备
    int deviceCount = 0;
    HIP_CHECK(hipGetDeviceCount(&deviceCount));
    if (deviceCount == 0) {
        std::cerr << "Error: No HIP devices found" << std::endl;
        return -1;
    }
    std::cout << "Found " << deviceCount << " HIP device(s) - Using device 0" << std::endl;
    HIP_CHECK(hipSetDevice(0));  // 设置当前线程使用的设备为0

    // 3. 执行所有测试模块：按功能分类执行，确保流程连贯
    test_graph_apis(release_resources);
    test_array_apis(release_resources);
    test_context_apis(release_resources);
    test_texture_apis(release_resources);
    test_user_object_apis(release_resources);
    test_memory_apis(release_resources);
    test_surface_apis(release_resources);
    test_surface_apis2(release_resources);
    test_texref_apis(release_resources);
    test_hipGreenCtxCreate_apis(release_resources);

    // 4. 测试总结：提示泄漏检测的预期结果
    std::cout << "\n====== Test Summary ======" << std::endl;
    if (release_resources) {
        std::cout << "All resources have been properly released." << std::endl;
        std::cout << "Expected Result: Memory leak tool reports NO LEAKS." << std::endl;
    } else {
        std::cout << "Some resources are intentionally not released (for leak test)." << std::endl;
        std::cout << "Expected Result: Memory leak tool detects leaks from the following APIs:" << std::endl;
        std::cout << "  - hipGraphCreate/hipGraphDestroy" << std::endl;
        std::cout << "  - hipArrayCreate/hipArrayDestroy (1D/3D/Mipmap)" << std::endl;
        std::cout << "  - hipCtxCreate/hipCtxDestroy (explicit context)" << std::endl;
        std::cout << "  - hipCreateTextureObject/hipDestroyTextureObject + 1D Array" << std::endl;
        std::cout << "  - hipUserObjectCreate/hipUserObjectRelease" << std::endl;
        std::cout << "  - hipMemCreate/hipMemRelease" << std::endl;
        std::cout << "  - hipMemPoolCreate/hipMemPoolDestroy" << std::endl;
        std::cout << "  - hipCreateSurfaceObject/hipDestroySurfaceObject + 2D Array" << std::endl;
        std::cout << "  - hipTexRefCreate/hipTexRefDestroy + bound 1D Array" << std::endl;
        std::cout << "  - hipGreenCtxCreate/hipGreenCtxDestroy + Stream" << std::endl;
    }

    std::cout << "\nTest completed successfully!" << std::endl;
    return 0;
}